Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@src/tl_templates/maca/barrier.h`:
- Around line 42-47: The inline asm in the synchronization primitives
(mbarrier_init, mbarrier_try_wait, mbarrier_wait, mbarrier_arrive (both
overloads), mbarrier_expect_tx, mbarrier_arrive_expect_tx,
mbarrier_cp_async_arrive, mbarrier_cp_async_arrive_noinc, fence_proxy_async,
fence_barrier_init) must declare the "memory" clobber so the compiler cannot
reorder loads/stores across these asm volatile blocks; update each asm
volatile(...) invocation to include "memory" in the clobber list (preserving
existing input/output constraints such as the "r" operands and any existing
clobbers) so the asm acts as a compiler-level memory barrier.

---

Nitpick comments:
In `@tilelang/tileop/gemm/gemm_maca_mma.py`:
- Around line 129-130: The two lambda assignments shared_layout_a and
shared_layout_b should be replaced with proper function definitions to satisfy
the linter (E731); define functions (e.g. def shared_layout_a(buf): return
make_maca_gemm_ab_layout(buf, 1 if self.trans_A else 2) and def
shared_layout_b(buf): return make_maca_gemm_ab_layout(buf, 2 if self.trans_B
else 1)) so callers remain unchanged and references to make_maca_gemm_ab_layout,
self.trans_A and self.trans_B are preserved.
- Around line 28-35: The integer conversion in _get_maca_gemm_k_pack currently
calls int(value) directly so non-numeric env values raise a generic ValueError;
wrap the conversion in a try/except around int(value) (or name it raw and then
parse) and on exception raise a clearer ValueError that includes the env var
name TILELANG_MACA_GEMM_K_PACK and the offending value (optionally include the
original exception message) before the existing k_pack < 1 check.

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@tilelang/intrinsics/maca_mma_macro_generator.py`:
- Around line 500-502: The ValueError in the else branch incorrectly states
"k_dim must be 0" while valid k_dim values handled are 4, 16, and 32; update the
raise in tilelang/intrinsics/maca_mma_macro_generator.py (the branch that
currently raises ValueError(f"k_dim must be 0 currently but got {k_dim}")) to a
clearer message that lists the allowed values and shows the actual value, e.g.
raise ValueError(f"k_dim must be one of {{4, 16, 32}} but got {k_dim}").
- Line 782: Remove the debugging print statement that emits kernel compilation
noise: delete the call to print(self.a_preshuffle) in the MacaMmaMacroGenerator
(or the method where self.a_preshuffle is used) so stdout is not polluted during
kernel compilation; ensure no other side-effects rely on that print and run
tests to confirm behavior unchanged.

In `@tilelang/intrinsics/mfma_layout.py`:
- Around line 113-128: The function shared_16x64_to_local_64x16_layout_B is a
copy of the _A variant and uses the wrong index mapping; replace its body so it
follows the established _B pattern: compute thread_id as j + (i // 16) * 16 and
local as i % 16, then return thread_id, local (update
shared_16x64_to_local_64x16_layout_B accordingly to match the intended mapping
and be consistent with shared_16x64_to_local_64x16_layout_A and other *_B
variants).

In `@tilelang/intrinsics/mfma_macro_generator.py`:
- Line 835: Remove the leftover debug print that emits self.a_preshuffle in
mfma_macro_generator.py: locate the print(self.a_preshuffle) call (inside the
MFMA macro generation code where self.a_preshuffle is referenced) and delete it;
if silent diagnostics are needed instead, replace it with a debug-level logging
call (e.g., using the module/class logger) rather than printing to stdout.

In `@tilelang/intrinsics/tcgen05_macro_generator.py`:
- Around line 22-27: The SwizzleMode IntEnum currently has incorrect numeric
mappings; update the SwizzleMode enum in tcgen05_macro_generator.py so the
members match the TCGen05 PTX spec: set NONE = 0, SWIZZLE_32B = 1, SWIZZLE_64B =
2, and SWIZZLE_128B = 3 (these correspond to bits 61–63 of the shared-memory
descriptor). Modify the class SwizzleMode to use these exact integer values so
code referencing SwizzleMode uses the correct PTX swizzle-mode encodings.

In `@tilelang/intrinsics/utils.py`:
- Around line 96-116: The function get_mma_micro_size currently types its dtype
as Literal["float16", "int8"] but the implementation and docstring also accept
"float8_e4m3" and "float8_e5m2"; update the type hint on get_mma_micro_size to
include those FP8 literals (e.g.,
Literal["float16","int8","float8_e4m3","float8_e5m2"]) or broaden to
str/Union[...] so the annotation matches the handled values, and ensure the
docstring stays consistent with the new annotation.

---

Outside diff comments:
In `@tilelang/maca/op/gemm/gemm_mma.py`:
- Around line 306-318: The accumulator clear call is inside the ki loop in
_gemm_srr which resets C_buf every iteration and discards partial products; move
the conditional clear_accum / T.clear(C_buf) out of the for ki in T.serial(0,
(block_K // macro_size_k)) loop so C_buf is cleared once before the loop begins
(preserve the existing clear_accum boolean check), leaving the subsequent
mma_emitter.ldmatrix_a(...) and mma_emitter.mma(...) calls unchanged.

---

Nitpick comments:
In `@tilelang/intrinsics/mma_sp_layout.py`:
- Around line 147-154: The dtype type hint on get_ldmatrix_offset_b is incorrect
for the implemented logic: update the annotation for the dtype parameter in
get_ldmatrix_offset_b to include "float32" (e.g.,
Literal["float16","int8","float32"]) or replace the Literal with a broader type
(like str or an enum) so it matches the branch that checks for dtype_bits == 32
and the float32 handling in the function body.

In `@tilelang/intrinsics/tcgen05_macro_generator.py`:
- Around line 278-304: The duplicated helper access_ptr_from (used in
tcgen05mma_ss and tcgen05mma_ts) should be extracted to a single module-level
function (or a shared class method) so both generators call the same
implementation; create a top-level def access_ptr_from(buffer_or_load_or_region,
access_type="r") that preserves the existing logic (Buffer, BufferLoad handling
with offset/stride computation, BufferRegion handling, and the same error
raising), replace the duplicated blocks in tcgen05mma_ss and tcgen05mma_ts with
calls to this new function, and ensure any local names (Buffer, BufferLoad,
BufferRegion, tvm) remain in scope or are imported so behavior is unchanged.

In `@tilelang/intrinsics/wgmma_macro_generator.py`:
- Around line 500-502: Replace the redundant construction of is_sr_conditions
and the any() call with a direct boolean based on transposed: the current block
that creates is_sr_conditions = [False]; is_sr_conditions.append(not
transposed); is_sr_axis_order = any(is_sr_conditions) should be simplified to
set is_sr_axis_order directly from not transposed (i.e., is_sr_axis_order = not
transposed) so remove the temporary list and any() usage; update any dependent
logic/comments in the function wgmma_macro_generator.py to reflect the
simplified variable.
- Around line 165-176: Rename the misspelled method `_determinate_swizzle_mode`
to `_determine_swizzle_mode` and update every call site that references
`_determinate_swizzle_mode` to use `_determine_swizzle_mode` instead;
specifically change the function definition and all invocations (e.g., the calls
currently referencing `_determinate_swizzle_mode` around where swizzle mode is
resolved) to the corrected name so references (and any imports/uses within the
same module) remain consistent and tests/imports continue to work.

In `@tilelang/intrinsics/wmma_macro_generator.py`:
- Around line 196-212: The macro parameter A_shared_buf in _warp_ldmatrix_a is
unused (the body reads A_buf from outer scope); either remove A_shared_buf from
the macro signature and from its invocation (return line) or change the macro
body to consistently use A_shared_buf instead of A_buf; apply the same fix for
the symmetric macro _warp_ldmatrix_b (remove unused B_shared_buf or switch uses
to B_shared_buf) and update any calls to _warp_ldmatrix_a/_warp_ldmatrix_b
accordingly so signatures match their invocations.

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@src/target/codegen_maca.cc`:
- Around line 1666-1685: The code emits PrintExpr(op->args[2]) as a template
parameter for tl::cp_async_gs, which can produce invalid C++ if the size isn't a
compile-time IntImm; update the ptx_cp_async handling (the block guarded by
op->op.same_as(builtin::ptx_cp_async())) to first normalize/validate op->args[2]
as an IntImm using the same helper used elsewhere (call
GetTileLangCPAsyncTransferBytes on op->args[2] to obtain a validated integer
byte count), then use that integer value (as a literal string) when emitting
both tl::cp_async_gs<...>(...) and tl::cp_async_gs_conditional<...>(...), and
emit a clear error/check if GetTileLangCPAsyncTransferBytes fails; apply the
same pattern consistently wherever ptx_cp_async is handled (e.g., the other
backends).

In `@src/tl_templates/maca/barrier.h`:
- Around line 68-75: The inline assembly in the __forceinline__ barrier helper
defines non-unique labels LAB_WAIT and DONE which will clash when the function
is instantiated in multiple translation units; update the labels inside the asm
block (the ones referenced by the branch instructions in the mbarrier loop) to
use unique local labels such as LAB_WAIT_%= and DONE_%= (or switch to numeric
local labels like 1: with 1b/1f references) so each inlined instance gets
distinct label names and avoids duplicate-label assembler errors.

In `@src/tl_templates/maca/gemm_wsm.h`:
- Around line 76-83: The code always seeds C_f32 from accum, ignoring the
template flag clear_accum; change the logic in the gemm_ss_wsm routine so that
before the current loop that assigns C_f32 from accum you check clear_accum and
if true initialize C_f32 to zero (or skip loading accum entirely), otherwise
proceed to load accum into C_f32 as currently done; reference the C_f32 array,
accum pointer, and the clear_accum template parameter to implement the
conditional initialization so first-use outputs are not contaminated by stale
fragments.

In `@src/tl_templates/maca/gemm.h`:
- Around line 96-100: The remove_swizzle helper is checking sizeof(A_type) but
is used for the B fragment (e.g., called with tCrB.layout()), which strips the
wrong layout when A and B have different widths; change the constexpr condition
to inspect sizeof(B_type) instead of sizeof(A_type) so that
ComposedLayout<Args...> const &layout returns layout.layout_b() only when B_type
is 2 bytes (and otherwise returns layout), updating the remove_swizzle
implementation accordingly to reference B_type for the swizzle decision so
gemm() receives the correct B layout.

In `@testing/maca/language/test_tilelang_language_access_ptr_codegen.py`:
- Around line 168-169: The two MACA tests (e.g., the test function
test_maca_bsm_intrinsics_codegen and the similar test at lines ~206-207) are
unconditionally running; gate them with the project’s MACA availability marker
so they are skipped when MACA isn’t present. Locate the test definitions
(test_maca_bsm_intrinsics_codegen and the other MACA-targeted test) and add the
same pytest marker used elsewhere in this file (e.g., the MACA availability
skip/marker) so the tests are automatically skipped when the project’s MACA
flag/marker indicates MACA is unavailable.

In `@tilelang/contrib/mxcc.py`:
- Around line 92-94: Handle malformed TILELANG_MXCC_FLAGS by catching
shlex.split ValueError around the call that appends to cmd: when reading
extra_env_flags and calling shlex.split(extra_env_flags) (the code that mutates
cmd using variables cmd and extra_env_flags), wrap the split in try/except
ValueError and raise a clear, user-facing error (or exit with a descriptive
message) that includes the original exception text and the offending
extra_env_flags value so users see "malformed TILELANG_MXCC_FLAGS: <details>"
instead of a raw traceback.

In `@tilelang/language/gemm_op.py`:
- Around line 159-160: The gemm function added a new public parameter
annotations but the docstring Args: section for gemm does not mention it; update
the gemm docstring to document the annotations parameter (type, purpose,
default/optional behavior) alongside the existing arguments—reference the gemm
function and the annotations parameter in your description so users know what
values (e.g., dict | None) are expected and how annotations affect behavior;
keep the style consistent with the other Args entries in the existing docstring.

In `@tilelang/tileop/gemm/gemm_base.py`:
- Around line 168-170: The annotations property can return None and cause
callers like self.annotations.get(...) to crash; change the property
(annotations) to always return a dict by retrieving getattr(self.gemm_node,
"annotations", None) and normalizing it to an empty dict when falsy (or
converting to dict if needed) so callers can safely call .get() on the result.

---

Nitpick comments:
In `@tilelang/contrib/mxcc.py`:
- Around line 20-44: The docstring for compile_maca is missing documentation for
the TILELANG_MXCC_FLAGS environment variable; update the compile_maca function
docstring to mention that additional MXCC compiler flags can be supplied via the
TILELANG_MXCC_FLAGS environment variable (e.g., a space-separated string or
list-equivalent), describe its effect on the options passed to mxcc, and include
an example or note about precedence/format so users can discover and use this
feature.

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@src/tl_templates/maca/gemm_wsm.h`:
- Around line 11-143: The template instantiation silently permits unsupported
configurations; add compile-time guards inside gemm_ss_wsm to fail fast:
static_assert that trans_A==false and trans_B==false, num_warp_m==1 and
num_warp_n==1 (the hardwired warp layout), kPack==8 (used as K/8),
AStrideElements%8==0 (used to compute lda_vec), and that the implementation
expects a 4-stage schedule (Stage==4) / 4x4 accumulators (i.e., require any
template flags you depend on to match the hardcoded schedule); reference the
function name gemm_ss_wsm and the symbols ALdgOffset, BLdgOffset, WSM_Ldg,
lda_vec, K/8, and C_f32 when adding these static_asserts so the checks are
colocated with the hardcoded logic.

---

Nitpick comments:
In `@tilelang/maca/op/gemm/gemm_mma.py`:
- Around line 286-298: The WSM workspace allocation uses a magic constant
(0x8000) in _gemm_ss_wsm_template which should be clarified; replace the literal
with a named constant (e.g., WSM_SIZE_BYTES or WSM_WORKSPACE_SIZE = 0x8000) or
add an inline comment explaining “32KB required for WSM workspace per TL gemm
backend” so readers understand the size choice and update the T.alloc_shared
call to use that constant (refer to WSM and _gemm_ss_wsm_template to locate the
change).
- Around line 137-150: infer_layout calls _get_maca_gemm_k_pack(self.k_pack) and
lower calls it again, risking inconsistency if the env var changes; to fix, read
and cache k_pack once (e.g., store in self._cached_k_pack at construction or the
first call) and replace direct calls to _get_maca_gemm_k_pack(...) in both
infer_layout and lower with the cached value, or pass the resolved k_pack from
infer_layout into lower via the operator state so both use the identical k_pack;
update references in methods named infer_layout and lower and any helper uses
like _make_maca_gemm_emitter to consume the cached/passed k_pack.